<!-- HTML header for doxygen 1.8.3.1-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.4"/>
<title>CUB: Kernel I/O</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="extra_stylesheet.css" rel="stylesheet" type="text/css"/>
<link rel="shortcut icon" href="favicon.ico" type="image/x-icon" />
<script type="text/javascript">
  var _gaq = _gaq || [];
  _gaq.push(['_setAccount', 'UA-38890655-1']);
  _gaq.push(['_trackPageview']);
  (function() {
    var ga = document.createElement('script'); ga.type = 'text/javascript'; ga.async = true;
    ga.src = ('https:' == document.location.protocol ? 'https://ssl' : 'http://www') + '.google-analytics.com/ga.js';
    var s = document.getElementsByTagName('script')[0]; s.parentNode.insertBefore(ga, s);
  })();
</script>
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td style="padding-left: 0.5em;">
   <div id="projectname">CUB
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.4 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Main&#160;Page</span></a></li>
      <li><a href="modules.html"><span>Modules</span></a></li>
      <li><a href="annotated.html"><span>Classes</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
</div><!-- top -->
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Namespaces</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Typedefs</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(7)"><span class="SelectionMark">&#160;</span>Enumerations</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(8)"><span class="SelectionMark">&#160;</span>Enumerator</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(9)"><span class="SelectionMark">&#160;</span>Groups</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div class="header">
  <div class="summary">
<a href="#enum-members">Enumerations</a>  </div>
  <div class="headertitle">
<div class="title">Kernel I/O</div>  </div>
</div><!--header-->
<div class="contents">
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="enum-members"></a>
Enumerations</h2></td></tr>
<tr class="memitem:ga023420f30fec7d4b187fc98f4fd2a55d"><td class="memItemLeft" align="right" valign="top">enum &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d">cub::PtxLoadModifier</a> { <br/>
&#160;&#160;<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55daee682a9988eceaefd9a53a7ef92fe6e2">cub::LOAD_DEFAULT</a>, 
<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55dacec87ab52dfb122d1555bfdfe44c23ae">cub::LOAD_CA</a>, 
<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55dafad1bebebe5fe5ee6f6dfb5dc220009f">cub::LOAD_CG</a>, 
<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55da6f0ae5826ad5eee54046731f7c71c874">cub::LOAD_CS</a>, 
<br/>
&#160;&#160;<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55da83161be6f5901c6559583c0f7ea96076">cub::LOAD_CV</a>, 
<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55da0217758d7be24e5fd450d1d49aec36b8">cub::LOAD_LDG</a>, 
<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55da390f7351eaf6b5b6ea23b4ec089a8236">cub::LOAD_VOLATILE</a>
<br/>
 }</td></tr>
<tr class="memdesc:ga023420f30fec7d4b187fc98f4fd2a55d"><td class="mdescLeft">&#160;</td><td class="mdescRight">Enumeration of PTX cache-modifiers for memory load operations.  <a href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d">More...</a><br/></td></tr>
<tr class="separator:ga023420f30fec7d4b187fc98f4fd2a55d"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gae9c7d6a6af7104f528509182ac9c9da2"><td class="memItemLeft" align="right" valign="top">enum &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2">cub::PtxStoreModifier</a> { <br/>
&#160;&#160;<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2a434aa4b3efc8e0c0ce4f1a00cdad26bd">cub::STORE_DEFAULT</a>, 
<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2afd29f1e82f0a9dc8258fb7977b0a8237">cub::STORE_WB</a>, 
<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2aacea07ea298b89dd1962a40b4823652d">cub::STORE_CG</a>, 
<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2a00ae8891d1acad179d134fdd60d7839b">cub::STORE_CS</a>, 
<br/>
&#160;&#160;<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2a5bf33b361cb9f3229588c95402fc484f">cub::STORE_WT</a>, 
<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2aba537dcdd7a709bcdff169e014dcaf08">cub::STORE_VOLATILE</a>
<br/>
 }</td></tr>
<tr class="memdesc:gae9c7d6a6af7104f528509182ac9c9da2"><td class="mdescLeft">&#160;</td><td class="mdescRight">Enumeration of PTX cache-modifiers for memory store operations.  <a href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2">More...</a><br/></td></tr>
<tr class="separator:gae9c7d6a6af7104f528509182ac9c9da2"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="member-group"></a>
Simple I/O</h2></td></tr>
<tr class="memitem:ga2467f2337aeeea3a3a1b39d19996d3c2"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga2467f2337aeeea3a3a1b39d19996d3c2"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ <br class="typebreak"/>
std::iterator_traits<br class="typebreak"/>
&lt; InputIteratorRA &gt;<br class="typebreak"/>
::value_type&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga2467f2337aeeea3a3a1b39d19996d3c2">cub::ThreadLoad</a> (InputIteratorRA itr)</td></tr>
<tr class="memdesc:ga2467f2337aeeea3a3a1b39d19996d3c2"><td class="mdescLeft">&#160;</td><td class="mdescRight">Thread utility for reading memory using <a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifiers.  <a href="#ga2467f2337aeeea3a3a1b39d19996d3c2">More...</a><br/></td></tr>
<tr class="separator:ga2467f2337aeeea3a3a1b39d19996d3c2"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga870b5b156d3d0a65a8c551884b18a2d6"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename OutputIteratorRA , typename T &gt; </td></tr>
<tr class="memitem:ga870b5b156d3d0a65a8c551884b18a2d6"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga870b5b156d3d0a65a8c551884b18a2d6">cub::ThreadStore</a> (OutputIteratorRA itr, T val)</td></tr>
<tr class="memdesc:ga870b5b156d3d0a65a8c551884b18a2d6"><td class="mdescLeft">&#160;</td><td class="mdescRight">Thread utility for writing memory using <a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2" title="Enumeration of PTX cache-modifiers for memory store operations. ">cub::PtxStoreModifier</a> cache modifiers.  <a href="#ga870b5b156d3d0a65a8c551884b18a2d6">More...</a><br/></td></tr>
<tr class="separator:ga870b5b156d3d0a65a8c551884b18a2d6"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="member-group"></a>
Blocked I/O</h2></td></tr>
<tr class="memitem:ga0f245535ce8c023d6d85bcfc0bf9907e"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga0f245535ce8c023d6d85bcfc0bf9907e"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga0f245535ce8c023d6d85bcfc0bf9907e">cub::LoadBlocked</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:ga0f245535ce8c023d6d85bcfc0bf9907e"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier.  <a href="#ga0f245535ce8c023d6d85bcfc0bf9907e">More...</a><br/></td></tr>
<tr class="separator:ga0f245535ce8c023d6d85bcfc0bf9907e"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gae694d6fa9abb0ef67802d12d2c19516c"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:gae694d6fa9abb0ef67802d12d2c19516c"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gae694d6fa9abb0ef67802d12d2c19516c">cub::LoadBlocked</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:gae694d6fa9abb0ef67802d12d2c19516c"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier, guarded by range.  <a href="#gae694d6fa9abb0ef67802d12d2c19516c">More...</a><br/></td></tr>
<tr class="separator:gae694d6fa9abb0ef67802d12d2c19516c"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga4aa27e3c558ba1ba7c306c834ad67b0f"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga4aa27e3c558ba1ba7c306c834ad67b0f"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga4aa27e3c558ba1ba7c306c834ad67b0f">cub::LoadBlocked</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items, T oob_default)</td></tr>
<tr class="memdesc:ga4aa27e3c558ba1ba7c306c834ad67b0f"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements..  <a href="#ga4aa27e3c558ba1ba7c306c834ad67b0f">More...</a><br/></td></tr>
<tr class="separator:ga4aa27e3c558ba1ba7c306c834ad67b0f"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga0e69bdc5b96bad62c9501af85532fc7b"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:ga0e69bdc5b96bad62c9501af85532fc7b"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga0e69bdc5b96bad62c9501af85532fc7b">cub::StoreBlocked</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:ga0e69bdc5b96bad62c9501af85532fc7b"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a blocked arrangement of items across a thread block into a linear segment of items using the specified cache modifier.  <a href="#ga0e69bdc5b96bad62c9501af85532fc7b">More...</a><br/></td></tr>
<tr class="separator:ga0e69bdc5b96bad62c9501af85532fc7b"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gac84f07cf1f61e45b26290515e9b4d3ae"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:gac84f07cf1f61e45b26290515e9b4d3ae"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gac84f07cf1f61e45b26290515e9b4d3ae">cub::StoreBlocked</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:gac84f07cf1f61e45b26290515e9b4d3ae"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a blocked arrangement of items across a thread block into a linear segment of items using the specified cache modifier, guarded by range.  <a href="#gac84f07cf1f61e45b26290515e9b4d3ae">More...</a><br/></td></tr>
<tr class="separator:gac84f07cf1f61e45b26290515e9b4d3ae"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="member-group"></a>
Striped I/O</h2></td></tr>
<tr class="memitem:gaa32b5c960e511b67901f4a5ba423048e"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:gaa32b5c960e511b67901f4a5ba423048e"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gaa32b5c960e511b67901f4a5ba423048e">cub::LoadStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:gaa32b5c960e511b67901f4a5ba423048e"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a striped arrangement across the thread block using the specified cache modifier.  <a href="#gaa32b5c960e511b67901f4a5ba423048e">More...</a><br/></td></tr>
<tr class="separator:gaa32b5c960e511b67901f4a5ba423048e"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga5a9c12543d3c373375c0bf03e9385a1f"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga5a9c12543d3c373375c0bf03e9385a1f"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga5a9c12543d3c373375c0bf03e9385a1f">cub::LoadStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:ga5a9c12543d3c373375c0bf03e9385a1f"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a striped arrangement across the thread block using the specified cache modifier, guarded by range.  <a href="#ga5a9c12543d3c373375c0bf03e9385a1f">More...</a><br/></td></tr>
<tr class="separator:ga5a9c12543d3c373375c0bf03e9385a1f"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga020063c6e3b8ffb09a87a00fe6823381"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga020063c6e3b8ffb09a87a00fe6823381"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga020063c6e3b8ffb09a87a00fe6823381">cub::LoadStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items, T oob_default)</td></tr>
<tr class="memdesc:ga020063c6e3b8ffb09a87a00fe6823381"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a striped arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements.  <a href="#ga020063c6e3b8ffb09a87a00fe6823381">More...</a><br/></td></tr>
<tr class="separator:ga020063c6e3b8ffb09a87a00fe6823381"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gadce0282ac2ac679aa9d046e534da6af2"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:gadce0282ac2ac679aa9d046e534da6af2"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gadce0282ac2ac679aa9d046e534da6af2">cub::StoreStriped</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:gadce0282ac2ac679aa9d046e534da6af2"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier.  <a href="#gadce0282ac2ac679aa9d046e534da6af2">More...</a><br/></td></tr>
<tr class="separator:gadce0282ac2ac679aa9d046e534da6af2"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga5d9cc4924d2bd640be9cb86e4754ce8c"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:ga5d9cc4924d2bd640be9cb86e4754ce8c"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga5d9cc4924d2bd640be9cb86e4754ce8c">cub::StoreStriped</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:ga5d9cc4924d2bd640be9cb86e4754ce8c"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier, guarded by range.  <a href="#ga5d9cc4924d2bd640be9cb86e4754ce8c">More...</a><br/></td></tr>
<tr class="separator:ga5d9cc4924d2bd640be9cb86e4754ce8c"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="member-group"></a>
Warp-striped I/O</h2></td></tr>
<tr class="memitem:gaea39c52d05d5f8d078ff76b154fad26d"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:gaea39c52d05d5f8d078ff76b154fad26d"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gaea39c52d05d5f8d078ff76b154fad26d">cub::LoadWarpStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:gaea39c52d05d5f8d078ff76b154fad26d"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a warp-striped arrangement across the thread block using the specified cache modifier.  <a href="#gaea39c52d05d5f8d078ff76b154fad26d">More...</a><br/></td></tr>
<tr class="separator:gaea39c52d05d5f8d078ff76b154fad26d"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga4486da2d836785d42f48ae379e8dd85a"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga4486da2d836785d42f48ae379e8dd85a"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga4486da2d836785d42f48ae379e8dd85a">cub::LoadWarpStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:ga4486da2d836785d42f48ae379e8dd85a"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a warp-striped arrangement across the thread block using the specified cache modifier, guarded by range.  <a href="#ga4486da2d836785d42f48ae379e8dd85a">More...</a><br/></td></tr>
<tr class="separator:ga4486da2d836785d42f48ae379e8dd85a"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga92a5618a110d7963bbc6f54637b904b1"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga92a5618a110d7963bbc6f54637b904b1"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga92a5618a110d7963bbc6f54637b904b1">cub::LoadWarpStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items, T oob_default)</td></tr>
<tr class="memdesc:ga92a5618a110d7963bbc6f54637b904b1"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a warp-striped arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements.  <a href="#ga92a5618a110d7963bbc6f54637b904b1">More...</a><br/></td></tr>
<tr class="separator:ga92a5618a110d7963bbc6f54637b904b1"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga07c4d38e195128da47fb9a8c582072d7"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:ga07c4d38e195128da47fb9a8c582072d7"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga07c4d38e195128da47fb9a8c582072d7">cub::StoreWarpStriped</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:ga07c4d38e195128da47fb9a8c582072d7"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a warp-striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier.  <a href="#ga07c4d38e195128da47fb9a8c582072d7">More...</a><br/></td></tr>
<tr class="separator:ga07c4d38e195128da47fb9a8c582072d7"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga7b3a3b0c99fbfb1afc76b5fde54ba3dc"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:ga7b3a3b0c99fbfb1afc76b5fde54ba3dc"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga7b3a3b0c99fbfb1afc76b5fde54ba3dc">cub::StoreWarpStriped</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:ga7b3a3b0c99fbfb1afc76b5fde54ba3dc"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a warp-striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier, guarded by range.  <a href="#ga7b3a3b0c99fbfb1afc76b5fde54ba3dc">More...</a><br/></td></tr>
<tr class="separator:ga7b3a3b0c99fbfb1afc76b5fde54ba3dc"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="member-group"></a>
Blocked, vectorized I/O</h2></td></tr>
<tr class="memitem:gadeb7216e9de14b5bd22f952cd1775eda"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD&gt; </td></tr>
<tr class="memitem:gadeb7216e9de14b5bd22f952cd1775eda"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gadeb7216e9de14b5bd22f952cd1775eda">cub::LoadBlockedVectorized</a> (int linear_tid, T *block_ptr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:gadeb7216e9de14b5bd22f952cd1775eda"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier.  <a href="#gadeb7216e9de14b5bd22f952cd1775eda">More...</a><br/></td></tr>
<tr class="separator:gadeb7216e9de14b5bd22f952cd1775eda"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gac805e3fdfaecc3be16f81b0af5c1e4db"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD&gt; </td></tr>
<tr class="memitem:gac805e3fdfaecc3be16f81b0af5c1e4db"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gac805e3fdfaecc3be16f81b0af5c1e4db">cub::StoreBlockedVectorized</a> (int linear_tid, T *block_ptr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:gac805e3fdfaecc3be16f81b0af5c1e4db"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a blocked arrangement of items across a thread block into a linear segment of items using the specified cache modifier.  <a href="#gac805e3fdfaecc3be16f81b0af5c1e4db">More...</a><br/></td></tr>
<tr class="separator:gac805e3fdfaecc3be16f81b0af5c1e4db"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
<h2 class="groupheader">Enumeration Type Documentation</h2>
<a class="anchor" id="ga023420f30fec7d4b187fc98f4fd2a55d"></a>
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">enum <a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d">cub::PtxLoadModifier</a></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Enumeration of PTX cache-modifiers for memory load operations. </p>
<table class="fieldtable">
<tr><th colspan="2">Enumerator</th></tr><tr><td class="fieldname"><em><a class="anchor" id="gga023420f30fec7d4b187fc98f4fd2a55daee682a9988eceaefd9a53a7ef92fe6e2"></a>LOAD_DEFAULT</em>&nbsp;</td><td class="fielddoc">
<p>Default (no modifier) </p>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="gga023420f30fec7d4b187fc98f4fd2a55dacec87ab52dfb122d1555bfdfe44c23ae"></a>LOAD_CA</em>&nbsp;</td><td class="fielddoc">
<p>Cache at all levels. </p>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="gga023420f30fec7d4b187fc98f4fd2a55dafad1bebebe5fe5ee6f6dfb5dc220009f"></a>LOAD_CG</em>&nbsp;</td><td class="fielddoc">
<p>Cache at global level. </p>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="gga023420f30fec7d4b187fc98f4fd2a55da6f0ae5826ad5eee54046731f7c71c874"></a>LOAD_CS</em>&nbsp;</td><td class="fielddoc">
<p>Cache streaming (likely to be accessed once) </p>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="gga023420f30fec7d4b187fc98f4fd2a55da83161be6f5901c6559583c0f7ea96076"></a>LOAD_CV</em>&nbsp;</td><td class="fielddoc">
<p>Cache as volatile (including cached system lines) </p>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="gga023420f30fec7d4b187fc98f4fd2a55da0217758d7be24e5fd450d1d49aec36b8"></a>LOAD_LDG</em>&nbsp;</td><td class="fielddoc">
<p>Cache as texture. </p>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="gga023420f30fec7d4b187fc98f4fd2a55da390f7351eaf6b5b6ea23b4ec089a8236"></a>LOAD_VOLATILE</em>&nbsp;</td><td class="fielddoc">
<p>Volatile (any memory space) </p>
</td></tr>
</table>

<p>Definition at line <a class="el" href="thread__load_8cuh_source.html#l00062">62</a> of file <a class="el" href="thread__load_8cuh_source.html">thread_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="gae9c7d6a6af7104f528509182ac9c9da2"></a>
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">enum <a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2">cub::PtxStoreModifier</a></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Enumeration of PTX cache-modifiers for memory store operations. </p>
<table class="fieldtable">
<tr><th colspan="2">Enumerator</th></tr><tr><td class="fieldname"><em><a class="anchor" id="ggae9c7d6a6af7104f528509182ac9c9da2a434aa4b3efc8e0c0ce4f1a00cdad26bd"></a>STORE_DEFAULT</em>&nbsp;</td><td class="fielddoc">
<p>Default (no modifier) </p>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="ggae9c7d6a6af7104f528509182ac9c9da2afd29f1e82f0a9dc8258fb7977b0a8237"></a>STORE_WB</em>&nbsp;</td><td class="fielddoc">
<p>Cache write-back all coherent levels. </p>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="ggae9c7d6a6af7104f528509182ac9c9da2aacea07ea298b89dd1962a40b4823652d"></a>STORE_CG</em>&nbsp;</td><td class="fielddoc">
<p>Cache at global level. </p>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="ggae9c7d6a6af7104f528509182ac9c9da2a00ae8891d1acad179d134fdd60d7839b"></a>STORE_CS</em>&nbsp;</td><td class="fielddoc">
<p>Cache streaming (likely to be accessed once) </p>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="ggae9c7d6a6af7104f528509182ac9c9da2a5bf33b361cb9f3229588c95402fc484f"></a>STORE_WT</em>&nbsp;</td><td class="fielddoc">
<p>Cache write-through (to system memory) </p>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="ggae9c7d6a6af7104f528509182ac9c9da2aba537dcdd7a709bcdff169e014dcaf08"></a>STORE_VOLATILE</em>&nbsp;</td><td class="fielddoc">
<p>Volatile shared (any memory space) </p>
</td></tr>
</table>

<p>Definition at line <a class="el" href="thread__store_8cuh_source.html#l00061">61</a> of file <a class="el" href="thread__store_8cuh_source.html">thread_store.cuh</a>.</p>

</div>
</div>
<h2 class="groupheader">Function Documentation</h2>
<a class="anchor" id="ga0f245535ce8c023d6d85bcfc0bf9907e"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::LoadBlocked </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">InputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to load. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">InputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base input iterator for loading from </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">items</td><td>Data to load </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__load_8cuh_source.html#l00078">78</a> of file <a class="el" href="block__load_8cuh_source.html">block_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="gae694d6fa9abb0ef67802d12d2c19516c"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::LoadBlocked </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">InputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>valid_items</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier, guarded by range. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to load. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">InputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base input iterator for loading from </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">items</td><td>Data to load </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">valid_items</td><td>Number of valid items to load </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__load_8cuh_source.html#l00107">107</a> of file <a class="el" href="block__load_8cuh_source.html">block_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ga4aa27e3c558ba1ba7c306c834ad67b0f"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::LoadBlocked </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">InputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>valid_items</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>oob_default</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements.. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to load. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">InputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base input iterator for loading from </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">items</td><td>Data to load </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">valid_items</td><td>Number of valid items to load </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">oob_default</td><td>Default value to assign out-of-bound items </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__load_8cuh_source.html#l00141">141</a> of file <a class="el" href="block__load_8cuh_source.html">block_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="gadeb7216e9de14b5bd22f952cd1775eda"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD&gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::LoadBlockedVectorized </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T *&#160;</td>
          <td class="paramname"><em>block_ptr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>The input offset (<code>block_ptr</code> + <code>block_offset</code>) must be quad-item aligned</p>
<p>The following conditions will prevent vectorization and loading will fall back to <a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401a2d4d8900d7e697e9dac4062e97d3d835">cub::BLOCK_LOAD_DIRECT</a>:</p>
<ul>
<li><code>ITEMS_PER_THREAD</code> is odd</li>
<li>The data type <code>T</code> is not a built-in primitive or CUDA vector type (e.g., <code>short</code>, <code>int2</code>, <code>double</code>, <code>float2</code>, etc.)</li>
</ul>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to load. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_ptr</td><td>Input pointer for loading from </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">items</td><td>Data to load </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__load_8cuh_source.html#l00422">422</a> of file <a class="el" href="block__load_8cuh_source.html">block_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="gaa32b5c960e511b67901f4a5ba423048e"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxLoadModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::LoadStriped </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">InputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Load a linear segment of items into a striped arrangement across the thread block using the specified cache modifier. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>striped arrangement</em></a> of elements across threads, where the items owned by each thread have a ranked distance of <code>BLOCK_THREADS</code> between them.</p>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">BLOCK_THREADS</td><td>The thread block size in threads </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to load. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">InputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base input iterator for loading from </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">items</td><td>Data to load </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__load_8cuh_source.html#l00185">185</a> of file <a class="el" href="block__load_8cuh_source.html">block_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ga5a9c12543d3c373375c0bf03e9385a1f"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxLoadModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::LoadStriped </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">InputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>valid_items</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Load a linear segment of items into a striped arrangement across the thread block using the specified cache modifier, guarded by range. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>striped arrangement</em></a> of elements across threads, where the items owned by each thread have a ranked distance of <code>BLOCK_THREADS</code> between them.</p>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">BLOCK_THREADS</td><td>The thread block size in threads </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to load. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">InputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base input iterator for loading from </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">items</td><td>Data to load </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">valid_items</td><td>Number of valid items to load </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__load_8cuh_source.html#l00215">215</a> of file <a class="el" href="block__load_8cuh_source.html">block_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ga020063c6e3b8ffb09a87a00fe6823381"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxLoadModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::LoadStriped </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">InputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>valid_items</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>oob_default</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Load a linear segment of items into a striped arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>striped arrangement</em></a> of elements across threads, where the items owned by each thread have a ranked distance of <code>BLOCK_THREADS</code> between them.</p>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">BLOCK_THREADS</td><td>The thread block size in threads </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to load. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">InputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base input iterator for loading from </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">items</td><td>Data to load </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">valid_items</td><td>Number of valid items to load </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">oob_default</td><td>Default value to assign out-of-bound items </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__load_8cuh_source.html#l00251">251</a> of file <a class="el" href="block__load_8cuh_source.html">block_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="gaea39c52d05d5f8d078ff76b154fad26d"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::LoadWarpStriped </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">InputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Load a linear segment of items into a warp-striped arrangement across the thread block using the specified cache modifier. </p>
<p>Assumes a <em>warp-striped arrangement</em> of elements across threads, where each warp owns a segment of (<code>WARP_THREADS</code> * <code>ITEMS_PER_THREAD</code>) consecutively ranked items, and the items owned by each thread have a ranked distance of <code>WARP_THREADS</code> between them.</p>
<dl class="section user"><dt>Usage Considerations</dt><dd>The number of threads in the thread block must be a multiple of the architecture's warp size.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to load. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">InputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base input iterator for loading from </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">items</td><td>Data to load </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__load_8cuh_source.html#l00296">296</a> of file <a class="el" href="block__load_8cuh_source.html">block_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ga4486da2d836785d42f48ae379e8dd85a"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::LoadWarpStriped </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">InputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>valid_items</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Load a linear segment of items into a warp-striped arrangement across the thread block using the specified cache modifier, guarded by range. </p>
<p>Assumes a <em>warp-striped arrangement</em> of elements across threads, where each warp owns a segment of (<code>WARP_THREADS</code> * <code>ITEMS_PER_THREAD</code>) consecutively ranked items, and the items owned by each thread have a ranked distance of <code>WARP_THREADS</code> between them.</p>
<dl class="section user"><dt>Usage Considerations</dt><dd>The number of threads in the thread block must be a multiple of the architecture's warp size.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to load. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">InputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base input iterator for loading from </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">items</td><td>Data to load </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">valid_items</td><td>Number of valid items to load </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__load_8cuh_source.html#l00332">332</a> of file <a class="el" href="block__load_8cuh_source.html">block_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ga92a5618a110d7963bbc6f54637b904b1"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::LoadWarpStriped </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">InputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>valid_items</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>oob_default</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Load a linear segment of items into a warp-striped arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements. </p>
<p>Assumes a <em>warp-striped arrangement</em> of elements across threads, where each warp owns a segment of (<code>WARP_THREADS</code> * <code>ITEMS_PER_THREAD</code>) consecutively ranked items, and the items owned by each thread have a ranked distance of <code>WARP_THREADS</code> between them.</p>
<dl class="section user"><dt>Usage Considerations</dt><dd>The number of threads in the thread block must be a multiple of the architecture's warp size.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to load. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">InputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for input (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base input iterator for loading from </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">items</td><td>Data to load </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">valid_items</td><td>Number of valid items to load </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">oob_default</td><td>Default value to assign out-of-bound items </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__load_8cuh_source.html#l00373">373</a> of file <a class="el" href="block__load_8cuh_source.html">block_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ga0e69bdc5b96bad62c9501af85532fc7b"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::StoreBlocked </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">OutputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Store a blocked arrangement of items across a thread block into a linear segment of items using the specified cache modifier. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2" title="Enumeration of PTX cache-modifiers for memory store operations. ">cub::PtxStoreModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to store. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">OutputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for output (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base output iterator for storing to </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">items</td><td>Data to store </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__store_8cuh_source.html#l00077">77</a> of file <a class="el" href="block__store_8cuh_source.html">block_store.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="gac84f07cf1f61e45b26290515e9b4d3ae"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::StoreBlocked </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">OutputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>valid_items</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Store a blocked arrangement of items across a thread block into a linear segment of items using the specified cache modifier, guarded by range. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2" title="Enumeration of PTX cache-modifiers for memory store operations. ">cub::PtxStoreModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to store. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">OutputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for output (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base output iterator for storing to </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">items</td><td>Data to store </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">valid_items</td><td>Number of valid items to write </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__store_8cuh_source.html#l00106">106</a> of file <a class="el" href="block__store_8cuh_source.html">block_store.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="gac805e3fdfaecc3be16f81b0af5c1e4db"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD&gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::StoreBlockedVectorized </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T *&#160;</td>
          <td class="paramname"><em>block_ptr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Store a blocked arrangement of items across a thread block into a linear segment of items using the specified cache modifier. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of elements across threads, where thread<sub><em>i</em></sub> owns the <em>i</em><sup>th</sup> segment of consecutively ranked items.</p>
<p>The output offset (<code>block_ptr</code> + <code>block_offset</code>) must be quad-item aligned, which is the default starting offset returned by <code>cudaMalloc()</code> </p>
<dl class="section user"><dt></dt><dd>The following conditions will prevent vectorization and storing will fall back to <a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700a9b8dcc7b6b06bcfc24af4f499523b880">cub::BLOCK_STORE_DIRECT</a>:<ul>
<li><code>ITEMS_PER_THREAD</code> is odd</li>
<li>The data type <code>T</code> is not a built-in primitive or CUDA vector type (e.g., <code>short</code>, <code>int2</code>, <code>double</code>, <code>float2</code>, etc.)</li>
</ul>
</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2" title="Enumeration of PTX cache-modifiers for memory store operations. ">cub::PtxStoreModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to store. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_ptr</td><td>Input pointer for storing from </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">items</td><td>Data to store </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__store_8cuh_source.html#l00311">311</a> of file <a class="el" href="block__store_8cuh_source.html">block_store.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="gadce0282ac2ac679aa9d046e534da6af2"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxStoreModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::StoreStriped </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">OutputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Store a striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>striped arrangement</em></a> of elements across threads, where the items owned by each thread have a ranked distance of <code>BLOCK_THREADS</code> between them.</p>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2" title="Enumeration of PTX cache-modifiers for memory store operations. ">cub::PtxStoreModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">BLOCK_THREADS</td><td>The thread block size in threads </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to store. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">OutputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for output (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base output iterator for storing to </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">items</td><td>Data to store </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__store_8cuh_source.html#l00149">149</a> of file <a class="el" href="block__store_8cuh_source.html">block_store.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ga5d9cc4924d2bd640be9cb86e4754ce8c"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxStoreModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::StoreStriped </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">OutputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>valid_items</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Store a striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier, guarded by range. </p>
<p>Assumes a <a href="index.html#sec5sec4"><em>striped arrangement</em></a> of elements across threads, where the items owned by each thread have a ranked distance of <code>BLOCK_THREADS</code> between them.</p>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2" title="Enumeration of PTX cache-modifiers for memory store operations. ">cub::PtxStoreModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">BLOCK_THREADS</td><td>The thread block size in threads </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to store. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">OutputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for output (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base output iterator for storing to </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">items</td><td>Data to store </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">valid_items</td><td>Number of valid items to write </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__store_8cuh_source.html#l00180">180</a> of file <a class="el" href="block__store_8cuh_source.html">block_store.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ga07c4d38e195128da47fb9a8c582072d7"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::StoreWarpStriped </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">OutputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD]&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Store a warp-striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier. </p>
<p>Assumes a <em>warp-striped arrangement</em> of elements across threads, where each warp owns a segment of (<code>WARP_THREADS</code> * <code>ITEMS_PER_THREAD</code>) consecutively ranked items, and the items owned by each thread have a ranked distance of <code>WARP_THREADS</code> between them.</p>
<dl class="section user"><dt>Usage Considerations</dt><dd>The number of threads in the thread block must be a multiple of the architecture's warp size.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2" title="Enumeration of PTX cache-modifiers for memory store operations. ">cub::PtxStoreModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to store. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">OutputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for output (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base output iterator for storing to </td></tr>
    <tr><td class="paramdir">[out]</td><td class="paramname">items</td><td>Data to load </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__store_8cuh_source.html#l00224">224</a> of file <a class="el" href="block__store_8cuh_source.html">block_store.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ga7b3a3b0c99fbfb1afc76b5fde54ba3dc"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::StoreWarpStriped </td>
          <td>(</td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>linear_tid</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">OutputIteratorRA&#160;</td>
          <td class="paramname"><em>block_itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T(&amp;)&#160;</td>
          <td class="paramname"><em>items</em>[ITEMS_PER_THREAD], </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">int&#160;</td>
          <td class="paramname"><em>valid_items</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Store a warp-striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier, guarded by range. </p>
<p>Assumes a <em>warp-striped arrangement</em> of elements across threads, where each warp owns a segment of (<code>WARP_THREADS</code> * <code>ITEMS_PER_THREAD</code>) consecutively ranked items, and the items owned by each thread have a ranked distance of <code>WARP_THREADS</code> between them.</p>
<dl class="section user"><dt>Usage Considerations</dt><dd>The number of threads in the thread block must be a multiple of the architecture's warp size.</dd></dl>
<dl class="tparams"><dt>Template Parameters</dt><dd>
  <table class="tparams">
    <tr><td class="paramname">MODIFIER</td><td><a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2" title="Enumeration of PTX cache-modifiers for memory store operations. ">cub::PtxStoreModifier</a> cache modifier. </td></tr>
    <tr><td class="paramname">T</td><td><b>[inferred]</b> The data type to store. </td></tr>
    <tr><td class="paramname">ITEMS_PER_THREAD</td><td><b>[inferred]</b> The number of consecutive items partitioned onto each thread. </td></tr>
    <tr><td class="paramname">OutputIteratorRA</td><td><b>[inferred]</b> The random-access iterator type for output (may be a simple pointer type). </td></tr>
  </table>
  </dd>
</dl>
<dl class="params"><dt>Parameters</dt><dd>
  <table class="params">
    <tr><td class="paramdir">[in]</td><td class="paramname">linear_tid</td><td>A suitable 1D thread-identifier for the calling thread (e.g., <code>(threadIdx.y * blockDim.x) + linear_tid</code> for 2D thread blocks) </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">block_itr</td><td>The thread block's base output iterator for storing to </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">items</td><td>Data to store </td></tr>
    <tr><td class="paramdir">[in]</td><td class="paramname">valid_items</td><td>Number of valid items to write </td></tr>
  </table>
  </dd>
</dl>

<p>Definition at line <a class="el" href="block__store_8cuh_source.html#l00260">260</a> of file <a class="el" href="block__store_8cuh_source.html">block_store.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="ga2467f2337aeeea3a3a1b39d19996d3c2"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxLoadModifier MODIFIER, typename InputIteratorRA &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ std::iterator_traits&lt;InputIteratorRA&gt;::value_type cub::ThreadLoad </td>
          <td>(</td>
          <td class="paramtype">InputIteratorRA&#160;</td>
          <td class="paramname"><em>itr</em>)</td><td></td>
          <td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Thread utility for reading memory using <a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifiers. </p>
<p>Cache modifiers will only be effected for built-in types (i.e., C++ primitives and CUDA vector-types).</p>
<p>For example: </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// 32-bit load using cache-global modifier:</span></div>
<div class="line"><span class="keywordtype">int</span> *d_in;</div>
<div class="line"><span class="keywordtype">int</span> val = cub::ThreadLoad&lt;cub::LOAD_CA&gt;(d_in + threadIdx.x);</div>
<div class="line"></div>
<div class="line"><span class="comment">// 16-bit load using default modifier</span></div>
<div class="line"><span class="keywordtype">short</span> *d_in;</div>
<div class="line"><span class="keywordtype">short</span> val = cub::ThreadLoad&lt;cub::LOAD_DEFAULT&gt;(d_in + threadIdx.x);</div>
<div class="line"></div>
<div class="line"><span class="comment">// 256-bit load using cache-volatile modifier</span></div>
<div class="line">double4 *d_in;</div>
<div class="line">double4 val = cub::ThreadLoad&lt;cub::LOAD_CV&gt;(d_in + threadIdx.x);</div>
<div class="line"></div>
<div class="line"><span class="comment">// 96-bit load using default cache modifier (ignoring LOAD_CS)</span></div>
<div class="line"><span class="keyword">struct </span>TestFoo { <span class="keywordtype">bool</span> a; <span class="keywordtype">short</span> b; };</div>
<div class="line">TestFoo *d_struct;</div>
<div class="line">TestFoo val = cub::ThreadLoad&lt;cub::LOAD_CS&gt;(d_in + threadIdx.x);</div>
</div><!-- fragment --> </dd></dl>

</div>
</div>
<a class="anchor" id="ga870b5b156d3d0a65a8c551884b18a2d6"></a>
<div class="memitem">
<div class="memproto">
<div class="memtemplate">
template&lt;PtxStoreModifier MODIFIER, typename OutputIteratorRA , typename T &gt; </div>
      <table class="memname">
        <tr>
          <td class="memname">__device__ __forceinline__ void cub::ThreadStore </td>
          <td>(</td>
          <td class="paramtype">OutputIteratorRA&#160;</td>
          <td class="paramname"><em>itr</em>, </td>
        </tr>
        <tr>
          <td class="paramkey"></td>
          <td></td>
          <td class="paramtype">T&#160;</td>
          <td class="paramname"><em>val</em>&#160;</td>
        </tr>
        <tr>
          <td></td>
          <td>)</td>
          <td></td><td></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>Thread utility for writing memory using <a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2" title="Enumeration of PTX cache-modifiers for memory store operations. ">cub::PtxStoreModifier</a> cache modifiers. </p>
<p>Cache modifiers will only be effected for built-in types (i.e., C++ primitives and CUDA vector-types).</p>
<p>For example: </p>
<dl class="section user"><dt></dt><dd><div class="fragment"><div class="line"><span class="preprocessor">#include &lt;<a class="code" href="cub_8cuh.html">cub/cub.cuh</a>&gt;</span></div>
<div class="line"></div>
<div class="line"><span class="comment">// 32-bit store using cache-global modifier:</span></div>
<div class="line"><span class="keywordtype">int</span> *d_out;</div>
<div class="line"><span class="keywordtype">int</span> val;</div>
<div class="line">cub::ThreadStore&lt;cub::STORE_CG&gt;(d_out + threadIdx.x, val);</div>
<div class="line"></div>
<div class="line"><span class="comment">// 16-bit store using default modifier</span></div>
<div class="line"><span class="keywordtype">short</span> *d_out;</div>
<div class="line"><span class="keywordtype">short</span> val;</div>
<div class="line">cub::ThreadStore&lt;cub::STORE_DEFAULT&gt;(d_out + threadIdx.x, val);</div>
<div class="line"></div>
<div class="line"><span class="comment">// 256-bit store using write-through modifier</span></div>
<div class="line">double4 *d_out;</div>
<div class="line">double4 val;</div>
<div class="line">cub::ThreadStore&lt;cub::STORE_WT&gt;(d_out + threadIdx.x, val);</div>
<div class="line"></div>
<div class="line"><span class="comment">// 96-bit store using default cache modifier (ignoring STORE_CS)</span></div>
<div class="line"><span class="keyword">struct </span>TestFoo { <span class="keywordtype">bool</span> a; <span class="keywordtype">short</span> b; };</div>
<div class="line">TestFoo *d_struct;</div>
<div class="line">TestFoo val;</div>
<div class="line">cub::ThreadStore&lt;cub::STORE_CS&gt;(d_out + threadIdx.x, val);</div>
</div><!-- fragment --> </dd></dl>

</div>
</div>
</div><!-- contents -->
<!-- HTML footer for doxygen 1.8.3.1-->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated on Fri Aug 23 2013 17:31:11 for CUB by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.4
<br>
&copy; 2013 NVIDIA Corporation
</small></address>
</body>
</html>
